#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include "max_min_kernel.hpp"
#include "stdio.h"
#include "stdlib.h"
#include "string.h"

__device__ void max_minUcharOneInt(unsigned int dataIn, uchar2 * max_minOut)
{
	uchar * puchar = (uchar*)&dataIn;

	uchar tempmaxLow = (puchar[0] > puchar[1]) ? puchar[0] : puchar[1];
	uchar tempminLow = (puchar[0] > puchar[1]) ? puchar[1] : puchar[0];

	uchar tempmaxHigh = (puchar[2] > puchar[3]) ? puchar[2] : puchar[3];
	uchar tempminHigh = (puchar[2] > puchar[3]) ? puchar[3] : puchar[2];

	max_minOut->x = (tempminLow > tempminHigh) ? tempminHigh : tempminLow;
	max_minOut->y = (tempmaxLow > tempmaxHigh) ? tempmaxLow : tempmaxHigh;
}

__global__ void calc_max_min_kernal(unsigned int * dev_imageBuffIn, unsigned int numDatas, uchar2* dev_max_minOut)
{
	//__shared__ uchar max_curBlock[256];
	//__shared__ uchar min_curBlock[256];

	extern __shared__ uchar max_curBlock[];
	uchar * min_curBlock = max_curBlock + blockDim.x;

	int indexIn = blockDim.x*blockIdx.x + threadIdx.x;

	if (indexIn >= numDatas)
	{//indexIn >= numDatas and threadIdx.x < 255
		max_curBlock[threadIdx.x] = 0;
		min_curBlock[threadIdx.x] = 255;

	}
	else
	{
		int numThreads = gridDim.x*blockDim.x;

		uchar2 max_minOut = { 0, 0 };
		unsigned int tempData;

		for (; indexIn < numDatas; indexIn += numThreads)
		{
			tempData = dev_imageBuffIn[indexIn];
			max_minUcharOneInt(tempData, &max_minOut);
			max_curBlock[threadIdx.x] = max_minOut.y;
			min_curBlock[threadIdx.x] = max_minOut.x;
		}
	}
	__syncthreads();

	//calc the max and min between the 256 datas in the shared memory
	for (int i = blockDim.x / 2; i>0; i = i >> 1)
	{
		if (threadIdx.x < i)
		{
			if (max_curBlock[threadIdx.x]<max_curBlock[threadIdx.x + i])
			{
				max_curBlock[threadIdx.x] = max_curBlock[threadIdx.x + i];
			}

			if (min_curBlock[threadIdx.x] > min_curBlock[threadIdx.x + i])
			{
				min_curBlock[threadIdx.x] = min_curBlock[threadIdx.x + i];
			}
		}
		__syncthreads();

	}

	if (threadIdx.x == 0)
	{
		dev_max_minOut[blockIdx.x].x = min_curBlock[threadIdx.x];
		dev_max_minOut[blockIdx.x].y = max_curBlock[threadIdx.x];
	}
}

__global__ void mergeMaxMin(uchar2 *devMaxminArrayIn, uchar2 * devMaxMinOut)
{
	int index = blockDim.x*blockIdx.x + threadIdx.x;
	extern __shared__ uchar max_Array[];
	uchar * min_Array = max_Array + blockDim.x;

	/*if (index >= 256)
	{
	return;
	}*/
	max_Array[index] = devMaxminArrayIn[index].y;
	min_Array[index] = devMaxminArrayIn[index].x;
	__syncthreads();
	for (int i = blockDim.x / 2; i > 0; i = i >> 1)
	{
		if (threadIdx.x < i)
		{
			if (max_Array[threadIdx.x] < max_Array[threadIdx.x + i])
			{
				max_Array[threadIdx.x] = max_Array[threadIdx.x + i];
			}

			if (min_Array[threadIdx.x] > min_Array[threadIdx.x + i])
			{
				min_Array[threadIdx.x] = min_Array[threadIdx.x + i];
			}
		}
		__syncthreads();
	}

	if (threadIdx.x == 0)
	{
		devMaxMinOut->x = min_Array[0];
		devMaxMinOut->y = max_Array[0];
	}
}

int calc_max_min_1d(uchar * dev_imageBuffIn, unsigned int width, unsigned int height, uchar * devBuff, uchar2* max_minOut)
{
	uchar2 *tempMaxMinArr = (uchar2 *)devBuff;
	//dim3 blockSize(1, (height>256)?256:height);
	int width4 = width / 4;

	int blockSize = 256;
	int gridSize = 256;
	while (width4 < blockSize)
	{
		if ((blockSize >> 1) < width4)
		{
			break;
		}
		else
		{
			blockSize = blockSize >> 1;
		}
	}

	while (height < gridSize)
	{
		if ((gridSize >> 1) < height)
		{
			break;
		}
		else
		{
			gridSize = gridSize >> 1;
		}
	}

	//printf("gridSize=%d,blockSize=%d\n", gridSize, blockSize);
	calc_max_min_kernal << <gridSize, blockSize, blockSize*sizeof(uchar2) >> >((unsigned int*)dev_imageBuffIn, height*width / 4, tempMaxMinArr);

	uchar2 * tempMaxMin = (uchar2 *)(devBuff + 256 * sizeof(uchar2));

	mergeMaxMin << <1, gridSize, gridSize*sizeof(uchar2) >> >(tempMaxMinArr, tempMaxMin);

	cudaMemcpy(max_minOut, tempMaxMin, sizeof(uchar2), cudaMemcpyDeviceToHost);
	return 0;

}
